//#pragma OPENCL EXTENSION cl_amd_printf:enable
#ifdef cl_amd_fp64 
#pragma OPENCL EXTENSION cl_amd_fp64 : enable
#endif
#ifdef cl_khr_fp64
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
#endif
#define FMAD128(a, b) \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      a = b * a + b; \
      b = a * b + a; \
      

__kernel void
sum(__global double* const dA,
    __global double* dResult
    )

{
    const size_t bx = get_group_id(0);
    const size_t tx = get_local_id(0);
     const unsigned pIndex = 256 * bx + tx;
    double a = dA[tx];
    double b = (double)1.0; 

	for (unsigned i = 1; i <= 15; ++i) {
        FMAD128(a, b);
        FMAD128(a, b);
        FMAD128(a, b);
        FMAD128(a, b);
        FMAD128(a, b);
        FMAD128(a, b);
        FMAD128(a, b);
        FMAD128(a, b);
        FMAD128(a, b);
        FMAD128(a, b);
        FMAD128(a, b);
        FMAD128(a, b);
        FMAD128(a, b);
        FMAD128(a, b);
        FMAD128(a, b);
        FMAD128(a, b);
	}
     dResult[pIndex] = a + b;
}

